[New offload driver][llvm-link][SYCL] Use LTO instead of llvm-link to link device input bitcodes #13395
Conversation
Signed-off-by: Sudarsanam, Arvind <arvind.sudarsanam@intel.com>
asudarsa
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
| @@ -2,10 +2,9 @@ | |||
|
|
|||
| /// Check for list of commands for standalone clang-linker-wrapper run for sycl | |||
| // RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.o,libsycl-complex.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s | |||
| // CHK-CMDS: "{{.*}}llvm-link" [[INPUT:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings | |||
| // CHK-CMDS-NEXT: "{{.*}}clang-offload-bundler" -type=o -targets=sycl-spir64-unknown-unknown -input={{.*}}libsycl-crt.o -output=[[FIRSTUNBUNDLEDLIB:.*]].bc -unbundle -allow-missing-bundles | |||
| // CHK-CMDS: "{{.*}}clang-offload-bundler" -type=o -targets=sycl-spir64-unknown-unknown -input={{.*}}libsycl-crt.o -output=[[FIRSTUNBUNDLEDLIB:.*]].bc -unbundle -allow-missing-bundles | |||
There was a problem hiding this comment.
Looks like some tests are failing
There was a problem hiding this comment.
It seems like we need some upstream changes for SPIRV backend to be registered for both spir and spirv target triples. I have talked about this to the SPIR-V backend team and this is in their TODO list. I will update this PR with latest information.
Thanks
| @@ -40,6 +40,7 @@ def do_configure(args): | |||
| libdevice_dir = os.path.join(abs_src_dir, "libdevice") | |||
| fusion_dir = os.path.join(abs_src_dir, "sycl-fusion") | |||
| llvm_targets_to_build = args.host_target | |||
| llvm_experimental_targets_to_build = 'SPIRV' | |||
There was a problem hiding this comment.
i wonder if @bader has an opinion on the effect on CI/other users since basically everyone except this team isn't going to use this so its wasted time/cycles for them
There was a problem hiding this comment.
I would prefer SPIR-V backend to be enabled by the option, which is disabled by default to avoid the impact on other DPC++ compiler developers.
I'm okay with turning on this option in CI if some reasonable number of tests are passing already and we want to avoid regressions in these tests.
To be honest, I expected much more changes to make LTO framework usable for SYCL compilation flow. Specifically replacing SPIR-V translator with SPIR-V back-end should introduce some number of .
@VyacheslavLevytskyy, will it be useful for you if we enable built of SPIR-V target in our CI (e.g. gather SYCL tests pass rate on additional platforms)? NOTE: sycl branch is usually behind main branch (few days, but can be weeks).
Overall, SPIR-V backend today is relatively small, so the overhead it not significant.
There was a problem hiding this comment.
There was a problem hiding this comment.
@bader Indeed, it'd be helpful to gather info about other environments easier, even with a delay of days/weeks. This move would make sense one day, but only on condition that we introduced a couple of minor changes into the translation chain under a new option. These would include (1) diagnostics from spirv-val about produced SPIR-V output, and (2) verification that llvm-spirv -r is able to produce LLVM IR from the SPIR-V output. Until this is changed all intermediate crashes are hidden, and we see only final numbers of fail/success cases, without any tips about possible causes of fails.
| @@ -963,13 +963,6 @@ static Expected<StringRef> linkDevice(ArrayRef<StringRef> InputFiles, | |||
| SmallVector<StringRef, 16> InputFilesVec; | |||
| for (StringRef InputFile : InputFiles) | |||
| InputFilesVec.emplace_back(InputFile); | |||
| // First llvm-link step. | |||
| auto LinkedFile = sycl::linkDeviceInputFiles(InputFilesVec, Args); | |||
There was a problem hiding this comment.
Did you test a few basic examples to make sure the results execute successfully on a GPU?
There was a problem hiding this comment.
Not exactly. Let me do that. I will keep this PR in draft mode till I can get to verify this. Thanks
|
This pull request is stale because it has been open 180 days with no activity. Remove stale label or comment or this will be automatically closed in 30 days. |
When using new offload model for SYCL offload, clang-linker-wrapper gathers all bitcode files and uses llvm-link to link them together.
This PR has the following changes:
Thanks